#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#define ACCUM_FLT4 float4
#define FLT float
#define FLT2 float2
#define FLT3 float3
#define FLT4 float4
#define TO_FLT4 convert_float4
#define TO_ACCUM_TYPE convert_float4
#define TO_ACCUM_FLT convert_float
#define READ_IMAGE read_imagef
#define WRITE_IMAGE write_imagef
__constant sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
#define ACCUM_FLT float
constant ACCUM_FLT At[24] = {
        1.0000000000f, 1.0000000000f, 1.0000000000f, 1.0000000000f, 1.0000000000f, 0.0000000000f, 
        0.0000000000f, 0.7071067691f, -0.7071067691f, 1.4142135382f, -1.4142135382f, 0.0000000000f, 
        0.0000000000f, 0.4999999702f, 0.4999999702f, 1.9999998808f, 1.9999998808f, 0.0000000000f, 
        0.0000000000f, 0.3535533547f, -0.3535533547f, 2.8284268379f, -2.8284268379f, 1.0000000000f, 
};
__kernel void winograd_36_to_4x4_read(
__global float4* dst_tensor_buffer,
  __global float4* src_tensor_buffer,
  __read_only image2d_t at_tex2d,
  __read_only image2d_t biases_tex2d,
  int4 shared_int4_0,                   // 120, 480, 480, 64  1, 36, 14400, 256
  int4 shared_int4_1) {                 // 36, 14400, 0, 0    1, 480, 480, 256
  int tile_id = get_global_id(0);       // 0 ~ 14400
  int DST_Y = get_global_id(1);         // 4
  int DST_Z = get_global_id(2);         // 64
  int tile_x = (tile_id % shared_int4_0.x) * 4;         // (id % 120) * 4, 这个Tile在x方向的坐标
  int tile_y = (tile_id / shared_int4_0.x) * 4 + DST_Y; // (id / 120) * 4 + range(0, 4), 这个Tile在y方向坐标加偏移量
  // (tile_x >= args.dst_tensor.Width() || tile_y >= args.dst_tensor.Height() || DST_Z >= args.dst_tensor.Slices())
  if (tile_x >= shared_int4_0.y || tile_y >= shared_int4_0.z || DST_Z >= shared_int4_0.w) {
    return; 
  }
  ACCUM_FLT4 I0, I1, I2, I3, I4, I5;

  // 从At中读取第dst_y行
  ACCUM_FLT at_ar[6];
  ACCUM_FLT4 t00 = TO_ACCUM_TYPE(read_imagef(at_tex2d, smp_none, (int2)(DST_Y * 2 + 0, 0)));
  ACCUM_FLT4 t01 = TO_ACCUM_TYPE(read_imagef(at_tex2d, smp_none, (int2)(DST_Y * 2 + 1, 0)));
  at_ar[0] = t00.x;
  at_ar[1] = t00.y;
  at_ar[2] = t00.z;
  at_ar[3] = t00.w;
  at_ar[4] = t01.x;
  at_ar[5] = t01.y;
  {
    ACCUM_FLT at = at_ar[0];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (0)) * shared_int4_1.y + (tile_id))]; // (dst_slice * 36 + 0) * 14400 + tile_id
    I0 = at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (1)) * shared_int4_1.y + (tile_id))];
    I1 = at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (2)) * shared_int4_1.y + (tile_id))];
    I2 = at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (3)) * shared_int4_1.y + (tile_id))];
    I3 = at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (4)) * shared_int4_1.y + (tile_id))];
    I4 = at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (5)) * shared_int4_1.y + (tile_id))];
    I5 = at * src5;
  }
  {
    ACCUM_FLT at = at_ar[1];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (6)) * shared_int4_1.y + (tile_id))];
    I0 += at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (7)) * shared_int4_1.y + (tile_id))];
    I1 += at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (8)) * shared_int4_1.y + (tile_id))];
    I2 += at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (9)) * shared_int4_1.y + (tile_id))];
    I3 += at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (10)) * shared_int4_1.y + (tile_id))];
    I4 += at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (11)) * shared_int4_1.y + (tile_id))];
    I5 += at * src5;
  }
  {
    ACCUM_FLT at = at_ar[2];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (12)) * shared_int4_1.y + (tile_id))];
    I0 += at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (13)) * shared_int4_1.y + (tile_id))];
    I1 += at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (14)) * shared_int4_1.y + (tile_id))];
    I2 += at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (15)) * shared_int4_1.y + (tile_id))];
    I3 += at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (16)) * shared_int4_1.y + (tile_id))];
    I4 += at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (17)) * shared_int4_1.y + (tile_id))];
    I5 += at * src5;
  }
  {
    ACCUM_FLT at = at_ar[3];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (18)) * shared_int4_1.y + (tile_id))];
    I0 += at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (19)) * shared_int4_1.y + (tile_id))];
    I1 += at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (20)) * shared_int4_1.y + (tile_id))];
    I2 += at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (21)) * shared_int4_1.y + (tile_id))];
    I3 += at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (22)) * shared_int4_1.y + (tile_id))];
    I4 += at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (23)) * shared_int4_1.y + (tile_id))];
    I5 += at * src5;
  }
  {
    ACCUM_FLT at = at_ar[4];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (24)) * shared_int4_1.y + (tile_id))];
    I0 += at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (25)) * shared_int4_1.y + (tile_id))];
    I1 += at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (26)) * shared_int4_1.y + (tile_id))];
    I2 += at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (27)) * shared_int4_1.y + (tile_id))];
    I3 += at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (28)) * shared_int4_1.y + (tile_id))];
    I4 += at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (29)) * shared_int4_1.y + (tile_id))];
    I5 += at * src5;
  }
  {
    ACCUM_FLT at = at_ar[5];
    ACCUM_FLT4 src0 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (30)) * shared_int4_1.y + (tile_id))];
    I0 += at * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (31)) * shared_int4_1.y + (tile_id))];
    I1 += at * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (32)) * shared_int4_1.y + (tile_id))];
    I2 += at * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (33)) * shared_int4_1.y + (tile_id))];
    I3 += at * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (34)) * shared_int4_1.y + (tile_id))];
    I4 += at * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[(((DST_Z) * shared_int4_1.x + (35)) * shared_int4_1.y + (tile_id))];
    I5 += at * src5;
  }
  ACCUM_FLT4 t0 = I1 + I2;
  ACCUM_FLT4 t1 = I3 + I4;
  FLT4 bias_val = read_imagef(biases_tex2d, smp_none, (int2)(DST_Z, 0));
  {
    FLT4 r0 = TO_FLT4(I0 + t0 + t1) + bias_val;
    dst_tensor_buffer[(((DST_Z) * shared_int4_0.z + (tile_y)) * shared_int4_0.y + (tile_x))] = r0;
    tile_x++;
  }
  ACCUM_FLT4 t2 = I1 - I2;
  ACCUM_FLT4 t3 = I3 - I4;
  if (tile_x < shared_int4_0.y) { // 在最开始阶段已经判断过了，所以能执行到这里这些条件肯定是能满足的
    FLT4 r0 = TO_FLT4(t2 * At[7] + t3 * At[9]) + bias_val;
    dst_tensor_buffer[(((DST_Z) * shared_int4_0.z + (tile_y)) * shared_int4_0.y + (tile_x))] = r0;
    tile_x++;
  }
  if (tile_x < shared_int4_0.y) {
    FLT4 r0 = TO_FLT4(t0 * At[13] + t1 * At[15]) + bias_val;
    dst_tensor_buffer[(((DST_Z) * shared_int4_0.z + (tile_y)) * shared_int4_0.y + (tile_x))] = r0;
    tile_x++;
  }
  if (tile_x < shared_int4_0.y) {
    FLT4 r0 = TO_FLT4(t2 * At[19] + t3 * At[21] + I5) + bias_val;
    dst_tensor_buffer[(((DST_Z) * shared_int4_0.z + (tile_y)) * shared_int4_0.y + (tile_x))] = r0;
    tile_x++;
  }
}